// OpenCL Data Parallel Model (SIMD)
// Multiply Add: MADD d = (a x b) + c, equivalent to 2 flops, optional FMA.
// 32 Bit Floating Point Code.

#define madd(a, b, c) (a * b + c)

#define madd128(a, b) \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \
    a = madd(b, a, b); \
    b = madd(a, b, a); \

// 2 flops * 128 * 16 * 4 = 16384 flops
#define BENCHMARK \
    for (size_t i=0; i<4; i++) { \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
        madd128(a, b); \
    } \

__kernel void fp32(__global float * const param0, __global float *param1) {
    const size_t gid = get_global_id(0);
    float a = (float)(param0[gid]);
    float b = (float)(1.0e-6f);
    BENCHMARK
    param1[gid] = a + b; // 16384 + 1 = 16385 flops
}

__kernel void fp32_2(__global float * const param0, __global float *param1) {
    const size_t gid = get_global_id(0);
    float2 a = (float2)(param0[gid], param0[gid]);
    float2 b = (float2)(1.0e-6f, 2.0e-6f);
    BENCHMARK
    param1[gid] = (a.x + b.x) + (a.y + b.y);
}

__kernel void fp32_4(__global float * const param0, __global float *param1) {
    const size_t gid = get_global_id(0);
    float4 a = (float4)(param0[gid], param0[gid], param0[gid], param0[gid]);
    float4 b = (float4)(1.0e-6f, 2.0e-6f, 3.0e-6f, 4.0e-6f);
    BENCHMARK
    // Intel HD Graphics bug. Needs driver version 9.17.10.3071 or greater.
    param1[gid] = (a.x + b.x) + (a.y + b.y) + (a.z + b.z) + (a.w + b.w);
}

__kernel void fp32_8(__global float * const param0, __global float *param1) {
    const size_t gid = get_global_id(0);
    float8 a = (float8)(param0[gid], param0[gid], param0[gid], param0[gid],
                        param0[gid], param0[gid], param0[gid], param0[gid]);
    float8 b = (float8)(1.0e-6f, 2.0e-6f, 3.0e-6f, 4.0e-6f,
                        5.0e-6f, 6.0e-6f, 7.0e-6f, 8.0e-6f);
    BENCHMARK
    param1[gid] = (a.s0 + b.s0) + (a.s1 + b.s1) + (a.s2 + b.s2) + (a.s3 + b.s3) +
                  (a.s4 + b.s4) + (a.s5 + b.s5) + (a.s6 + b.s6) + (a.s7 + b.s7);
}

/*__kernel void fp32_16(__global float * const param0, __global float *param1) {
    const size_t gid = get_global_id(0);
    float16 a = (float16)(param0[gid], param0[gid], param0[gid], param0[gid],
                          param0[gid], param0[gid], param0[gid], param0[gid],
                          param0[gid], param0[gid], param0[gid], param0[gid],
                          param0[gid], param0[gid], param0[gid], param0[gid]);
    float16 b = (float16)( 1.0e-6f,  2.0e-6f,  3.0e-6f,  4.0e-6f,
                           5.0e-6f,  6.0e-6f,  7.0e-6f,  8.0e-6f,
                           9.0e-6f, 10.0e-6f, 11.0e-6f, 12.0e-6f,
                          13.0e-6f, 14.0e-6f, 15.0e-6f, 16.0e-6f);
    BENCHMARK
    param1[gid] = (a.s0 + b.s0) + (a.s1 + b.s1) + (a.s2 + b.s2) + (a.s3 + b.s3) +
                  (a.s4 + b.s4) + (a.s5 + b.s5) + (a.s6 + b.s6) + (a.s7 + b.s7) +
                  (a.s8 + b.s8) + (a.s9 + b.s9) + (a.sa + b.sa) + (a.sb + b.sb) +
                  (a.sc + b.sc) + (a.sd + b.sd) + (a.se + b.se) + (a.sf + b.sf);
}*/
